@robertmaynard Do you think these are "real issues" that we should fix for C++ standards compliance (like avoiding UB) or are they potentially compiler bugs? Do you see these issues with newer GCC versions or older CUDA versions on ARM?

@robertmaynard Do you think these are "real issues" that we should fix for C++ standards compliance (like avoiding UB) or are they potentially compiler bugs? Do you see these issues with newer GCC versions or older CUDA versions on ARM?

These are all real issues that aren't exposed by the x86 compiler. I am still digging into why, but here is the background on why these failures are valid:

From the CUDA programming guide 14.5.3.1. Device Memory Space Specifiers:

A __device__, __constant__, __managed__ or __shared__ variable definition cannot have a class type with a non-empty constructor or a non-empty destructor. A constructor for a class type is considered empty at a point in the translation unit, if it is either a trivial constructor or it satisfies all of the following conditions:

• The constructor function has been defined.
• The constructor function has no parameters, the initializer list is empty and the function body is an empty compound statement.
• Its class has no virtual functions, no virtual base classes and no non-static data member initializers.
• The default constructors of all base classes of its class can be considered empty.
• For all the nonstatic data members of its class that are of class type (or array thereof), the default constructors can be considered empty.

The simpliest way to understand the above is that the __shared__ type must pass the std::is_trivially_constructible check, if it doesn't it fails to meet the requirements and could generate the warning. As an example if we look at the block_desc_s specification we can see that it fails to meet the is_trivially_constructible requirements ( https://godbolt.org/z/75rv9jd3s )

struct block_desc_s {
  block_desc_s() {};

  explicit constexpr block_desc_s(size_t offset_, uint32_t size_,
                                  uint32_t row_offset_, uint32_t first_row_,
                                  uint32_t num_rows_)
      : offset(offset_), size(size_), row_offset(row_offset_),
        first_row(first_row_), num_rows(num_rows_) {}

  size_t offset;
  uint32_t size;
  uint32_t row_offset;
  uint32_t first_row;
  uint32_t num_rows;
};

/merge